// Copyright (c) Meta Platforms, Inc. and affiliates.

#include <stdlib.h>
#include <string.h>

#include "openzl/codecs/common/bitstream/ff_bitstream.h"
#include "openzl/codecs/common/copy.h"
#include "openzl/codecs/common/window.h"
#include "openzl/codecs/entropy/deprecated/common_entropy.h"
#include "openzl/codecs/rolz/decode_decoder.h"
#include "openzl/common/allocation.h"
#include "openzl/common/cursor.h"
#include "openzl/shared/mem.h"
#include "openzl/shared/portability.h"
#include "openzl/zl_errors.h"

#define kMinOffset 16
#define kTokenOFBits 2
#define kTokenLLBits 4
#define kTokenMLBits 5

#define kTokenLLMask ((1u << kTokenLLBits) - 1)
#define kTokenMLMask ((1u << kTokenMLBits) - 1)

#define kTokenLLCopy (1u << kTokenLLBits)
#define kTokenMLCopy (1u << kTokenMLBits)

static ZS_decoderCtx* ZS_decoderCtx_create(void)
{
    return (ZS_decoderCtx*)ZL_malloc(sizeof(ZS_decoderCtx));
}

static void ZS_decoderCtx_release(ZS_decoderCtx* ctx)
{
    ZL_free(ctx);
}

static void ZS_decoderCtx_reset(ZS_decoderCtx* ctx)
{
    (void)ctx;
}

#define kMaxNumClusters 256

static ZL_Report decodeCodes(uint8_t* codes, size_t dstCapacity, ZL_RC* src);

static ZL_Report decodeLiterals(uint8_t* lits, size_t litCapacity, ZL_RC* src)
{
    return decodeCodes(lits, litCapacity, src);
}

static ZL_Report isRaw(ZL_RC src)
{
    ZL_Report type = ZS_Entropy_getType(ZL_RC_ptr(&src), ZL_RC_avail(&src));
    ZL_RET_R_IF_ERR(type);
    return ZL_returnValue(ZL_RES_value(type) == ZS_Entropy_Type_raw);
}

static ZL_Report getDecodedSize(ZL_RC src, size_t elementSize)
{
    return ZS_Entropy_getDecodedSize(
            ZL_RC_ptr(&src), ZL_RC_avail(&src), elementSize);
}

static uint8_t const* getRawBuffer(ZL_RC* src, size_t elementSize)
{
    ZL_Report const raw = isRaw(*src);
    if (ZL_isError(raw) || !ZL_validResult(raw))
        return NULL;

    ZL_Report offset =
            ZS_Entropy_getHeaderSize(ZL_RC_ptr(src), ZL_RC_avail(src));
    if (ZL_isError(offset))
        return NULL;
    ZL_Report report = ZS_Entropy_getEncodedSize(
            ZL_RC_ptr(src), ZL_RC_avail(src), elementSize);
    if (ZL_isError(report) || ZL_validResult(report) > ZL_RC_avail(src))
        return NULL;
    uint8_t const* buffer = ZL_RC_ptr(src) + ZL_validResult(offset);
    ZL_RC_advance(src, ZL_validResult(report));
    return buffer;
}

static ZL_Report decodeCodes(uint8_t* codes, size_t numSequences, ZL_RC* src)
{
    ZS_Entropy_DecodeParameters params = {
        .allowedTypes = ZS_Entropy_TypeMask_all,
        .tableManager = NULL,
    };
    return ZS_Entropy_decode(codes, numSequences, src, 1, &params);
}

static ZL_Report decodeCodes16(uint16_t* codes, size_t numSequences, ZL_RC* src)
{
    ZS_Entropy_DecodeParameters params = {
        .allowedTypes = ZS_Entropy_TypeMask_all,
        .tableManager = NULL,
    };
    return ZS_Entropy_decode(codes, numSequences, src, 2, &params);
}

#if ZL_HAS_AVX2

static void dpr(char const* name, __m128i const vec32)
{
    size_t n = 4;
    uint32_t data[4];
    _mm_storeu_si128((__m128i*)(void*)data, vec32);
    fprintf(stderr, "%s = [", name);
    for (size_t i = 0; i < n; ++i) {
        fprintf(stderr, "%x", data[i]);
        if (i + 1 < n)
            fprintf(stderr, ", ");
    }
    fprintf(stderr, "]\n");
}

static void dpr16(char const* name, __m128i const vec32)
{
    size_t n = 8;
    uint16_t data[8];
    _mm_storeu_si128((__m128i*)(void*)data, vec32);
    fprintf(stderr, "%s = [", name);
    for (size_t i = 0; i < n; ++i) {
        fprintf(stderr, "%x", data[i]);
        if (i + 1 < n)
            fprintf(stderr, ", ");
    }
    fprintf(stderr, "]\n");
}

#    define _ 0xFF
// clang-format off
static uint8_t const ZL_ALIGNED(16) ZL_UNUSED shuffle[8][16] = {
  { 0x00, 0x01, 0x02, 0x03, 0x00, 0x01, 0x02, 0x03, 0x00, 0x01, 0x02, 0x03, _, _, _, _ }, // 000
  { 0x04, 0x05, 0x06, 0x07, 0x04, 0x05, 0x06, 0x07, 0x04, 0x05, 0x06, 0x07, _, _, _, _ }, // 001
  { 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x04, 0x05, 0x06, 0x07, _, _, _, _ }, // 010
  { 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x08, 0x09, 0x0a, 0x0b, _, _, _, _ }, // 011
  { 0x00, 0x01, 0x02, 0x03, 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, _, _, _, _ }, // 100
  { 0x04, 0x05, 0x06, 0x07, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, _, _, _, _ }, // 101
  { 0x00, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, _, _, _, _ }, // 110
  { 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f, _, _, _, _ }, // 111
};

#undef _
#define _ 9
static int32_t const shuffle7[128][8] __attribute((aligned(32),unused)) = {
  { 0, 0, 0, 0, 0, 0, 0, _ }, // 0000000
  { 1, 1, 1, 1, 1, 1, 1, _ }, // 1000000
  { 0, 1, 1, 1, 1, 1, 1, _ }, // 0100000
  { 1, 2, 2, 2, 2, 2, 2, _ }, // 1100000
  { 0, 0, 1, 1, 1, 1, 1, _ }, // 0010000
  { 1, 1, 2, 2, 2, 2, 2, _ }, // 1010000
  { 0, 1, 2, 2, 2, 2, 2, _ }, // 0110000
  { 1, 2, 3, 3, 3, 3, 3, _ }, // 1110000
  { 0, 0, 0, 1, 1, 1, 1, _ }, // 0001000
  { 1, 1, 1, 2, 2, 2, 2, _ }, // 1001000
  { 0, 1, 1, 2, 2, 2, 2, _ }, // 0101000
  { 1, 2, 2, 3, 3, 3, 3, _ }, // 1101000
  { 0, 0, 1, 2, 2, 2, 2, _ }, // 0011000
  { 1, 1, 2, 3, 3, 3, 3, _ }, // 1011000
  { 0, 1, 2, 3, 3, 3, 3, _ }, // 0111000
  { 1, 2, 3, 4, 4, 4, 4, _ }, // 1111000
  { 0, 0, 0, 0, 1, 1, 1, _ }, // 0000100
  { 1, 1, 1, 1, 2, 2, 2, _ }, // 1000100
  { 0, 1, 1, 1, 2, 2, 2, _ }, // 0100100
  { 1, 2, 2, 2, 3, 3, 3, _ }, // 1100100
  { 0, 0, 1, 1, 2, 2, 2, _ }, // 0010100
  { 1, 1, 2, 2, 3, 3, 3, _ }, // 1010100
  { 0, 1, 2, 2, 3, 3, 3, _ }, // 0110100
  { 1, 2, 3, 3, 4, 4, 4, _ }, // 1110100
  { 0, 0, 0, 1, 2, 2, 2, _ }, // 0001100
  { 1, 1, 1, 2, 3, 3, 3, _ }, // 1001100
  { 0, 1, 1, 2, 3, 3, 3, _ }, // 0101100
  { 1, 2, 2, 3, 4, 4, 4, _ }, // 1101100
  { 0, 0, 1, 2, 3, 3, 3, _ }, // 0011100
  { 1, 1, 2, 3, 4, 4, 4, _ }, // 1011100
  { 0, 1, 2, 3, 4, 4, 4, _ }, // 0111100
  { 1, 2, 3, 4, 5, 5, 5, _ }, // 1111100
  { 0, 0, 0, 0, 0, 1, 1, _ }, // 0000010
  { 1, 1, 1, 1, 1, 2, 2, _ }, // 1000010
  { 0, 1, 1, 1, 1, 2, 2, _ }, // 0100010
  { 1, 2, 2, 2, 2, 3, 3, _ }, // 1100010
  { 0, 0, 1, 1, 1, 2, 2, _ }, // 0010010
  { 1, 1, 2, 2, 2, 3, 3, _ }, // 1010010
  { 0, 1, 2, 2, 2, 3, 3, _ }, // 0110010
  { 1, 2, 3, 3, 3, 4, 4, _ }, // 1110010
  { 0, 0, 0, 1, 1, 2, 2, _ }, // 0001010
  { 1, 1, 1, 2, 2, 3, 3, _ }, // 1001010
  { 0, 1, 1, 2, 2, 3, 3, _ }, // 0101010
  { 1, 2, 2, 3, 3, 4, 4, _ }, // 1101010
  { 0, 0, 1, 2, 2, 3, 3, _ }, // 0011010
  { 1, 1, 2, 3, 3, 4, 4, _ }, // 1011010
  { 0, 1, 2, 3, 3, 4, 4, _ }, // 0111010
  { 1, 2, 3, 4, 4, 5, 5, _ }, // 1111010
  { 0, 0, 0, 0, 1, 2, 2, _ }, // 0000110
  { 1, 1, 1, 1, 2, 3, 3, _ }, // 1000110
  { 0, 1, 1, 1, 2, 3, 3, _ }, // 0100110
  { 1, 2, 2, 2, 3, 4, 4, _ }, // 1100110
  { 0, 0, 1, 1, 2, 3, 3, _ }, // 0010110
  { 1, 1, 2, 2, 3, 4, 4, _ }, // 1010110
  { 0, 1, 2, 2, 3, 4, 4, _ }, // 0110110
  { 1, 2, 3, 3, 4, 5, 5, _ }, // 1110110
  { 0, 0, 0, 1, 2, 3, 3, _ }, // 0001110
  { 1, 1, 1, 2, 3, 4, 4, _ }, // 1001110
  { 0, 1, 1, 2, 3, 4, 4, _ }, // 0101110
  { 1, 2, 2, 3, 4, 5, 5, _ }, // 1101110
  { 0, 0, 1, 2, 3, 4, 4, _ }, // 0011110
  { 1, 1, 2, 3, 4, 5, 5, _ }, // 1011110
  { 0, 1, 2, 3, 4, 5, 5, _ }, // 0111110
  { 1, 2, 3, 4, 5, 6, 6, _ }, // 1111110
  { 0, 0, 0, 0, 0, 0, 1, _ }, // 0000001
  { 1, 1, 1, 1, 1, 1, 2, _ }, // 1000001
  { 0, 1, 1, 1, 1, 1, 2, _ }, // 0100001
  { 1, 2, 2, 2, 2, 2, 3, _ }, // 1100001
  { 0, 0, 1, 1, 1, 1, 2, _ }, // 0010001
  { 1, 1, 2, 2, 2, 2, 3, _ }, // 1010001
  { 0, 1, 2, 2, 2, 2, 3, _ }, // 0110001
  { 1, 2, 3, 3, 3, 3, 4, _ }, // 1110001
  { 0, 0, 0, 1, 1, 1, 2, _ }, // 0001001
  { 1, 1, 1, 2, 2, 2, 3, _ }, // 1001001
  { 0, 1, 1, 2, 2, 2, 3, _ }, // 0101001
  { 1, 2, 2, 3, 3, 3, 4, _ }, // 1101001
  { 0, 0, 1, 2, 2, 2, 3, _ }, // 0011001
  { 1, 1, 2, 3, 3, 3, 4, _ }, // 1011001
  { 0, 1, 2, 3, 3, 3, 4, _ }, // 0111001
  { 1, 2, 3, 4, 4, 4, 5, _ }, // 1111001
  { 0, 0, 0, 0, 1, 1, 2, _ }, // 0000101
  { 1, 1, 1, 1, 2, 2, 3, _ }, // 1000101
  { 0, 1, 1, 1, 2, 2, 3, _ }, // 0100101
  { 1, 2, 2, 2, 3, 3, 4, _ }, // 1100101
  { 0, 0, 1, 1, 2, 2, 3, _ }, // 0010101
  { 1, 1, 2, 2, 3, 3, 4, _ }, // 1010101
  { 0, 1, 2, 2, 3, 3, 4, _ }, // 0110101
  { 1, 2, 3, 3, 4, 4, 5, _ }, // 1110101
  { 0, 0, 0, 1, 2, 2, 3, _ }, // 0001101
  { 1, 1, 1, 2, 3, 3, 4, _ }, // 1001101
  { 0, 1, 1, 2, 3, 3, 4, _ }, // 0101101
  { 1, 2, 2, 3, 4, 4, 5, _ }, // 1101101
  { 0, 0, 1, 2, 3, 3, 4, _ }, // 0011101
  { 1, 1, 2, 3, 4, 4, 5, _ }, // 1011101
  { 0, 1, 2, 3, 4, 4, 5, _ }, // 0111101
  { 1, 2, 3, 4, 5, 5, 6, _ }, // 1111101
  { 0, 0, 0, 0, 0, 1, 2, _ }, // 0000011
  { 1, 1, 1, 1, 1, 2, 3, _ }, // 1000011
  { 0, 1, 1, 1, 1, 2, 3, _ }, // 0100011
  { 1, 2, 2, 2, 2, 3, 4, _ }, // 1100011
  { 0, 0, 1, 1, 1, 2, 3, _ }, // 0010011
  { 1, 1, 2, 2, 2, 3, 4, _ }, // 1010011
  { 0, 1, 2, 2, 2, 3, 4, _ }, // 0110011
  { 1, 2, 3, 3, 3, 4, 5, _ }, // 1110011
  { 0, 0, 0, 1, 1, 2, 3, _ }, // 0001011
  { 1, 1, 1, 2, 2, 3, 4, _ }, // 1001011
  { 0, 1, 1, 2, 2, 3, 4, _ }, // 0101011
  { 1, 2, 2, 3, 3, 4, 5, _ }, // 1101011
  { 0, 0, 1, 2, 2, 3, 4, _ }, // 0011011
  { 1, 1, 2, 3, 3, 4, 5, _ }, // 1011011
  { 0, 1, 2, 3, 3, 4, 5, _ }, // 0111011
  { 1, 2, 3, 4, 4, 5, 6, _ }, // 1111011
  { 0, 0, 0, 0, 1, 2, 3, _ }, // 0000111
  { 1, 1, 1, 1, 2, 3, 4, _ }, // 1000111
  { 0, 1, 1, 1, 2, 3, 4, _ }, // 0100111
  { 1, 2, 2, 2, 3, 4, 5, _ }, // 1100111
  { 0, 0, 1, 1, 2, 3, 4, _ }, // 0010111
  { 1, 1, 2, 2, 3, 4, 5, _ }, // 1010111
  { 0, 1, 2, 2, 3, 4, 5, _ }, // 0110111
  { 1, 2, 3, 3, 4, 5, 6, _ }, // 1110111
  { 0, 0, 0, 1, 2, 3, 4, _ }, // 0001111
  { 1, 1, 1, 2, 3, 4, 5, _ }, // 1001111
  { 0, 1, 1, 2, 3, 4, 5, _ }, // 0101111
  { 1, 2, 2, 3, 4, 5, 6, _ }, // 1101111
  { 0, 0, 1, 2, 3, 4, 5, _ }, // 0011111
  { 1, 1, 2, 3, 4, 5, 6, _ }, // 1011111
  { 0, 1, 2, 3, 4, 5, 6, _ }, // 0111111
  { 1, 2, 3, 4, 5, 6, 7, _ }, // 1111111
};
static uint8_t const ZL_ALIGNED(16) shuffle7_16[128][16] = {
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, _ }, // 0000000
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, _ }, // 1000000
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, _ }, // 0100000
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 1100000
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, _ }, // 0010000
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 1010000
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0110000
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1110000
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, _ }, // 0001000
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 1001000
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0101000
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1101000
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0011000
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1011000
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 0111000
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, _ }, // 1111000
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, _ }, // 0000100
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 1000100
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0100100
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1100100
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0010100
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1010100
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 0110100
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, _ }, // 1110100
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, _ }, // 0001100
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 1001100
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 0101100
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, _ }, // 1101100
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, _ }, // 0011100
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, _ }, // 1011100
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, _ }, // 0111100
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, 0xa, 0xb, _ }, // 1111100
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, _ }, // 0000010
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, _ }, // 1000010
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, _ }, // 0100010
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 1100010
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, _ }, // 0010010
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 1010010
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0110010
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1110010
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, _ }, // 0001010
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 1001010
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0101010
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1101010
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0011010
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1011010
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 0111010
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, _ }, // 1111010
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, _ }, // 0000110
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 1000110
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0100110
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1100110
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0010110
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1010110
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 0110110
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, _ }, // 1110110
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, _ }, // 0001110
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 1001110
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 0101110
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, _ }, // 1101110
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, _ }, // 0011110
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, _ }, // 1011110
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, _ }, // 0111110
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, 0xc, 0xd, _ }, // 1111110
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, _ }, // 0000001
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, _ }, // 1000001
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, _ }, // 0100001
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 1100001
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, _ }, // 0010001
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 1010001
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0110001
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1110001
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, _ }, // 0001001
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 1001001
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0101001
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1101001
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0011001
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1011001
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 0111001
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, _ }, // 1111001
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, _ }, // 0000101
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 1000101
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0100101
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1100101
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0010101
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1010101
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 0110101
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, _ }, // 1110101
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, _ }, // 0001101
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 1001101
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 0101101
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, _ }, // 1101101
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, _ }, // 0011101
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, _ }, // 1011101
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, _ }, // 0111101
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xa, 0xb, 0xc, 0xd, _ }, // 1111101
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, _ }, // 0000011
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, _ }, // 1000011
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, _ }, // 0100011
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 1100011
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, _ }, // 0010011
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 1010011
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0110011
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1110011
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, _ }, // 0001011
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 1001011
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0101011
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1101011
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0011011
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1011011
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 0111011
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, _ }, // 1111011
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, _ }, // 0000111
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 1000111
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0100111
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1100111
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0010111
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1010111
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 0110111
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, _ }, // 1110111
  { 0x0, 0x1, 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, _ }, // 0001111
  { 0x2, 0x3, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 1001111
  { 0x0, 0x1, 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 0101111
  { 0x2, 0x3, 0x4, 0x5, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, _ }, // 1101111
  { 0x0, 0x1, 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, _ }, // 0011111
  { 0x2, 0x3, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, _ }, // 1011111
  { 0x0, 0x1, 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, _ }, // 0111111
  { 0x2, 0x3, 0x4, 0x5, 0x6, 0x7, 0x8, 0x9, 0xa, 0xb, 0xc, 0xd, 0xe, 0xf, _ }, // 1111111
};

static uint8_t const ZL_ALIGNED(16) shuffle256_16[256][16] = {
    { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f }, // 00000000
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff }, // 01000000
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff }, // 10000000
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff }, // 11000000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff }, // 00010000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 01010000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 10010000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 11010000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff }, // 00100000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 01100000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 10100000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 11100000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff }, // 00110000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff }, // 01110000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff }, // 10110000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff }, // 11110000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff }, // 00000100
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 01000100
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 10000100
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 11000100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 00010100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 01010100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 10010100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 11010100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 00100100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 01100100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 10100100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 11100100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 00110100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 01110100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 10110100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 11110100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff }, // 00001000
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 01001000
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 10001000
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 11001000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 00011000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 01011000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 10011000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 11011000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 00101000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 01101000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 10101000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 11101000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 00111000
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 01111000
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 10111000
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 11111000
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff }, // 00001100
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff }, // 01001100
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff }, // 10001100
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff }, // 11001100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff }, // 00011100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff }, // 01011100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff }, // 10011100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x04, 0x05, 0x06, 0xff }, // 11011100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff }, // 00101100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff }, // 01101100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x04, 0x05, 0x06, 0xff }, // 10101100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff, 0x05, 0x06, 0x07, 0xff }, // 11101100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff }, // 00111100
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff, 0x04, 0x05, 0x06, 0xff }, // 01111100
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff, 0x05, 0x06, 0x07, 0xff }, // 10111100
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff, 0x06, 0x07, 0x08, 0xff }, // 11111100
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff }, // 00000001
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 01000001
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 10000001
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff }, // 11000001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 00010001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 01010001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 10010001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 11010001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 00100001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 01100001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 10100001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 11100001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff }, // 00110001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff }, // 01110001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff }, // 10110001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff }, // 11110001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff }, // 00000101
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 01000101
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 10000101
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 11000101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 00010101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 01010101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 10010101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 11010101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 00100101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 01100101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 10100101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 11100101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 00110101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 01110101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 10110101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff, 0x07, 0xff, 0xff, 0xff }, // 11110101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff }, // 00001001
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 01001001
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 10001001
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 11001001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff }, // 00011001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 01011001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 10011001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 11011001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff }, // 00101001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 01101001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 10101001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0xff, 0xff, 0xff }, // 11101001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff }, // 00111001
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0xff, 0xff, 0xff }, // 01111001
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0xff, 0xff, 0xff }, // 10111001
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff, 0x08, 0xff, 0xff, 0xff }, // 11111001
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff }, // 00001101
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff }, // 01001101
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff }, // 10001101
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff }, // 11001101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff }, // 00011101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff }, // 01011101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff }, // 10011101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0xff, 0xff, 0xff }, // 11011101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff }, // 00101101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff }, // 01101101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0xff, 0xff, 0xff }, // 10101101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0xff, 0xff, 0xff }, // 11101101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff }, // 00111101
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0xff, 0xff, 0xff }, // 01111101
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0xff, 0xff, 0xff }, // 10111101
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff, 0x09, 0xff, 0xff, 0xff }, // 11111101
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff }, // 00000010
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 01000010
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 10000010
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff }, // 11000010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 00010010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 01010010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 10010010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 11010010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 00100010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 01100010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 10100010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 11100010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff }, // 00110010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff }, // 01110010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff }, // 10110010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff }, // 11110010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff }, // 00000110
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 01000110
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 10000110
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 11000110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 00010110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 01010110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 10010110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 11010110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 00100110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 01100110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 10100110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 11100110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 00110110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 01110110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 10110110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff, 0x07, 0x08, 0xff, 0xff }, // 11110110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff }, // 00001010
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 01001010
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 10001010
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 11001010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff }, // 00011010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 01011010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 10011010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 11011010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff }, // 00101010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 01101010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 10101010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0x08, 0xff, 0xff }, // 11101010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff }, // 00111010
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0xff, 0xff }, // 01111010
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0x08, 0xff, 0xff }, // 10111010
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff, 0x08, 0x09, 0xff, 0xff }, // 11111010
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff }, // 00001110
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff }, // 01001110
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff }, // 10001110
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff }, // 11001110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff }, // 00011110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff }, // 01011110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff }, // 10011110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0xff, 0xff }, // 11011110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff }, // 00101110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff }, // 01101110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0xff, 0xff }, // 10101110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0x09, 0xff, 0xff }, // 11101110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff }, // 00111110
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0xff, 0xff }, // 01111110
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0x09, 0xff, 0xff }, // 10111110
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff, 0x09, 0x0a, 0xff, 0xff }, // 11111110
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff }, // 00000011
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff }, // 01000011
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 10000011
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff }, // 11000011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff }, // 00010011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 01010011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 10010011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 11010011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 00100011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 01100011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 10100011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 11100011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff }, // 00110011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff }, // 01110011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff }, // 10110011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff }, // 11110011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff }, // 00000111
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 01000111
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 10000111
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 11000111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 00010111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 01010111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 10010111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 11010111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 00100111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 01100111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 10100111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0x07, 0x08, 0xff }, // 11100111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 00110111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0xff, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 01110111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0xff, 0xff, 0xff, 0x06, 0x07, 0x08, 0xff }, // 10110111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0xff, 0xff, 0xff, 0x07, 0x08, 0x09, 0xff }, // 11110111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff }, // 00001011
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 01001011
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 10001011
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 11001011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff }, // 00011011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 01011011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 10011011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0x08, 0xff }, // 11011011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff }, // 00101011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 01101011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0x08, 0xff }, // 10101011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0x08, 0x09, 0xff }, // 11101011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff }, // 00111011
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0xff, 0xff, 0x06, 0x07, 0x08, 0xff }, // 01111011
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0xff, 0xff, 0x07, 0x08, 0x09, 0xff }, // 10111011
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0xff, 0xff, 0x08, 0x09, 0x0a, 0xff }, // 11111011
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff }, // 00001111
  { 0x00, 0xff, 0xff, 0xff, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff }, // 01001111
  { 0x00, 0x01, 0xff, 0xff, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff }, // 10001111
  { 0x00, 0x01, 0x02, 0xff, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff }, // 11001111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff }, // 00011111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0xff, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff }, // 01011111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0xff, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff }, // 10011111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0xff, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0x09, 0xff }, // 11011111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff }, // 00101111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0xff, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff }, // 01101111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0xff, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0x09, 0xff }, // 10101111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0xff, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0x09, 0x0a, 0xff }, // 11101111
  { 0x0c, 0x0d, 0x0e, 0x0f, 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff }, // 00111111
  { 0x00, 0xff, 0xff, 0xff, 0x01, 0x02, 0x03, 0xff, 0x04, 0x05, 0x06, 0xff, 0x07, 0x08, 0x09, 0xff }, // 01111111
  { 0x00, 0x01, 0xff, 0xff, 0x02, 0x03, 0x04, 0xff, 0x05, 0x06, 0x07, 0xff, 0x08, 0x09, 0x0a, 0xff }, // 10111111
  { 0x00, 0x01, 0x02, 0xff, 0x03, 0x04, 0x05, 0xff, 0x06, 0x07, 0x08, 0xff, 0x09, 0x0a, 0x0b, 0xff }, // 11111111
};
static uint8_t const consumed256_16[256] = {
     0, // 00000000
   1, // 01000000
   2, // 10000000
   3, // 11000000
   1, // 00010000
   2, // 01010000
   3, // 10010000
   4, // 11010000
   2, // 00100000
   3, // 01100000
   4, // 10100000
   5, // 11100000
   3, // 00110000
   4, // 01110000
   5, // 10110000
   6, // 11110000
   1, // 00000100
   2, // 01000100
   3, // 10000100
   4, // 11000100
   2, // 00010100
   3, // 01010100
   4, // 10010100
   5, // 11010100
   3, // 00100100
   4, // 01100100
   5, // 10100100
   6, // 11100100
   4, // 00110100
   5, // 01110100
   6, // 10110100
   7, // 11110100
   2, // 00001000
   3, // 01001000
   4, // 10001000
   5, // 11001000
   3, // 00011000
   4, // 01011000
   5, // 10011000
   6, // 11011000
   4, // 00101000
   5, // 01101000
   6, // 10101000
   7, // 11101000
   5, // 00111000
   6, // 01111000
   7, // 10111000
   8, // 11111000
   3, // 00001100
   4, // 01001100
   5, // 10001100
   6, // 11001100
   4, // 00011100
   5, // 01011100
   6, // 10011100
   7, // 11011100
   5, // 00101100
   6, // 01101100
   7, // 10101100
   8, // 11101100
   6, // 00111100
   7, // 01111100
   8, // 10111100
   9, // 11111100
   1, // 00000001
   2, // 01000001
   3, // 10000001
   4, // 11000001
   2, // 00010001
   3, // 01010001
   4, // 10010001
   5, // 11010001
   3, // 00100001
   4, // 01100001
   5, // 10100001
   6, // 11100001
   4, // 00110001
   5, // 01110001
   6, // 10110001
   7, // 11110001
   2, // 00000101
   3, // 01000101
   4, // 10000101
   5, // 11000101
   3, // 00010101
   4, // 01010101
   5, // 10010101
   6, // 11010101
   4, // 00100101
   5, // 01100101
   6, // 10100101
   7, // 11100101
   5, // 00110101
   6, // 01110101
   7, // 10110101
   8, // 11110101
   3, // 00001001
   4, // 01001001
   5, // 10001001
   6, // 11001001
   4, // 00011001
   5, // 01011001
   6, // 10011001
   7, // 11011001
   5, // 00101001
   6, // 01101001
   7, // 10101001
   8, // 11101001
   6, // 00111001
   7, // 01111001
   8, // 10111001
   9, // 11111001
   4, // 00001101
   5, // 01001101
   6, // 10001101
   7, // 11001101
   5, // 00011101
   6, // 01011101
   7, // 10011101
   8, // 11011101
   6, // 00101101
   7, // 01101101
   8, // 10101101
   9, // 11101101
   7, // 00111101
   8, // 01111101
   9, // 10111101
  10, // 11111101
   2, // 00000010
   3, // 01000010
   4, // 10000010
   5, // 11000010
   3, // 00010010
   4, // 01010010
   5, // 10010010
   6, // 11010010
   4, // 00100010
   5, // 01100010
   6, // 10100010
   7, // 11100010
   5, // 00110010
   6, // 01110010
   7, // 10110010
   8, // 11110010
   3, // 00000110
   4, // 01000110
   5, // 10000110
   6, // 11000110
   4, // 00010110
   5, // 01010110
   6, // 10010110
   7, // 11010110
   5, // 00100110
   6, // 01100110
   7, // 10100110
   8, // 11100110
   6, // 00110110
   7, // 01110110
   8, // 10110110
   9, // 11110110
   4, // 00001010
   5, // 01001010
   6, // 10001010
   7, // 11001010
   5, // 00011010
   6, // 01011010
   7, // 10011010
   8, // 11011010
   6, // 00101010
   7, // 01101010
   8, // 10101010
   9, // 11101010
   7, // 00111010
   8, // 01111010
   9, // 10111010
  10, // 11111010
   5, // 00001110
   6, // 01001110
   7, // 10001110
   8, // 11001110
   6, // 00011110
   7, // 01011110
   8, // 10011110
   9, // 11011110
   7, // 00101110
   8, // 01101110
   9, // 10101110
  10, // 11101110
   8, // 00111110
   9, // 01111110
  10, // 10111110
  11, // 11111110
   3, // 00000011
   4, // 01000011
   5, // 10000011
   6, // 11000011
   4, // 00010011
   5, // 01010011
   6, // 10010011
   7, // 11010011
   5, // 00100011
   6, // 01100011
   7, // 10100011
   8, // 11100011
   6, // 00110011
   7, // 01110011
   8, // 10110011
   9, // 11110011
   4, // 00000111
   5, // 01000111
   6, // 10000111
   7, // 11000111
   5, // 00010111
   6, // 01010111
   7, // 10010111
   8, // 11010111
   6, // 00100111
   7, // 01100111
   8, // 10100111
   9, // 11100111
   7, // 00110111
   8, // 01110111
   9, // 10110111
  10, // 11110111
   5, // 00001011
   6, // 01001011
   7, // 10001011
   8, // 11001011
   6, // 00011011
   7, // 01011011
   8, // 10011011
   9, // 11011011
   7, // 00101011
   8, // 01101011
   9, // 10101011
  10, // 11101011
   8, // 00111011
   9, // 01111011
  10, // 10111011
  11, // 11111011
   6, // 00001111
   7, // 01001111
   8, // 10001111
   9, // 11001111
   7, // 00011111
   8, // 01011111
   9, // 10011111
  10, // 11011111
   8, // 00101111
   9, // 01101111
  10, // 10101111
  11, // 11101111
   9, // 00111111
  10, // 01111111
  11, // 10111111
  12, // 11111111
};
// clang-format on

ZL_FORCE_INLINE uint32_t const*
getOffsets3(uint32_t out[4], uint32_t const* offsets, uint16_t const* tokens)
{
    uint64_t const mask = _pext_u64(ZL_readLE64(tokens), 0x000100010001ULL);
    // fprintf(stderr, "mask = 0x%x\n", ZL_readLE32(tokens) & 0x00010101);
    __m128i const shuffleV =
            _mm_load_si128((__m128i const*)(void const*)shuffle[mask]);
    __m128i offsetsV = _mm_loadu_si128((__m128i const*)(void const*)offsets);
    // dpr("input", offsetsV);
    offsetsV = _mm_shuffle_epi8(offsetsV, shuffleV);
    // dpr("shuff", shuffleV);
    // dpr("shuff", offsetsV);
    _mm_storeu_si128((__m128i*)(void*)out, offsetsV);
    return offsets + _mm_popcnt_u64(mask);
}

ZL_FORCE_INLINE uint8_t const* getOffsets4(
        uint32_t out[4],
        uint8_t const* offsets,
        uint16_t const* tokens,
        uint32_t rep)
{
    uint64_t const mask = _pext_u64(ZL_readLE64(tokens), 0x0003000300030003ULL);
    ZL_ASSERT_LT(mask, 256);
    // fprintf(stderr, "mask = 0x%x\n", ZL_readLE32(tokens) & 0x00010101);
    __m128i const shuffleV =
            _mm_load_si128((__m128i const*)(void const*)shuffle256_16[mask]);
    size_t const consumed = consumed256_16[mask];
    __m128i offsetsV = _mm_loadu_si128((__m128i const*)(void const*)offsets);
    offsetsV         = _mm_insert_epi32(offsetsV, (int)rep, 3);
    // dpr("input", offsetsV);
    offsetsV = _mm_shuffle_epi8(offsetsV, shuffleV);
    // dpr("shuff", shuffleV);
    // dpr("shuff", offsetsV);
    _mm_storeu_si128((__m128i*)(void*)out, offsetsV);
    return offsets + consumed;
}

// ZL_FORCE_INLINE ZL_TARGET_ATTRIBUTE("avx2,bmi2") uint32_t const* getOffsets7(
//     uint32_t out[8],
//     uint32_t const* offsets,
//     uint8_t const* tokens) {
//   uint64_t const mask = _pext_u64(ZL_readLE64(tokens),
//   0x0001010101010101ULL);
//   // fprintf(stderr, "mask = 0x%x\n", ZL_readLE32(tokens) & 0x00010101);
//   __m256i const shuffleV =
//       _mm256_load_si256((__m256i const*)(void const*)shuffle7[mask]);
//   __m256i offsetsV = _mm256_loadu_si256((__m256i const*)(void
//   const*)offsets);
//   // dpr("input", offsetsV);
//   offsetsV = _mm256_permutevar8x32_epi32(offsetsV, shuffleV);
//   // dpr("shuff", shuffleV);
//   // dpr("shuff", offsetsV);
//   _mm256_storeu_si256((__m256i*)(void*)out, offsetsV);
//   return offsets + _mm_popcnt_u64(mask);
// }

#    define getOffsets getOffsets3

ZL_FORCE_INLINE uint8_t const*
getOffsets16(uint16_t out[16], uint8_t const* offsets, uint8_t const* tokens)
{
    uint64_t const mask = _pext_u64(ZL_readLE64(tokens), 0x0001010101010101ULL);
    // fprintf(stderr, "mask = 0x%x\n", ZL_readLE32(tokens) & 0x00010101);
    __m128i const shuffleV =
            _mm_load_si128((__m128i const*)(void const*)shuffle7_16[mask]);
    __m128i offsetsV = _mm_loadu_si128((__m128i const*)(void const*)offsets);
    // dpr16("input", offsetsV);
    offsetsV = _mm_shuffle_epi8(offsetsV, shuffleV);
    // dpr16("shuff", shuffleV);
    // dpr16("shuff", offsetsV);
    _mm_storeu_si128((__m128i*)(void*)out, offsetsV);
    return offsets + 2 * _mm_popcnt_u64(mask);
}

ZL_FORCE_INLINE bool isFast(uint16_t const* tokens)
{
    // return (tokens[0] & ~1) && (tokens[1] & ~1) && (tokens[2] & ~1);
    __m256i tokensV = _mm256_loadu_si256((__m256i const*)(void const*)tokens);
    __m256i const maskV = _mm256_set1_epi16((int16_t)0xFFFC);
    tokensV             = _mm256_and_si256(tokensV, maskV);
    tokensV             = _mm256_cmpeq_epi16(tokensV, _mm256_setzero_si256());
    return _mm256_movemask_epi8(tokensV) == 0;
}

#else // ZL_HAS_AVX2

ZL_FORCE_INLINE uint8_t const* getOffsets4(
        uint32_t out[4],
        uint8_t const* offsets,
        uint16_t const* tokens,
        uint32_t rep)
{
    for (size_t i = 0; i < 4; ++i) {
        int const ocode = tokens[i] & 3;
        if (ocode == 0) {
            out[i] = rep;
        } else if (ocode == 1) {
            out[i] = *offsets++;
            rep    = out[i];
        } else if (ocode == 2) {
            out[i] = ZL_readLE16(offsets);
            offsets += 2;
            rep = out[i];
        } else {
            out[i] = ZL_readLE24(offsets);
            offsets += 3;
            rep = out[i];
        }
    }
    return offsets;
}

ZL_FORCE_INLINE bool isFast(uint16_t const* tokens)
{
    bool fast = true;
    for (size_t i = 0; i < 16; ++i) {
        fast &= ((tokens[i] & ~3) != 0);
    }
    return fast;
}

#endif

static ZL_Report getBitstream(ZS_BitDStreamFF* bitstream, ZL_RC* in)
{
    ZL_TRY_LET_CONST_T(uint64_t, offbitsSize, ZL_RC_popVarint(in));
    uint8_t const* const offbitsBuffer = ZL_RC_ptr(in);
    ZL_RET_R_IF_GT(srcSize_tooSmall, offbitsSize, ZL_RC_avail(in));
    ZL_RC_advance(in, offbitsSize);
    ZS_BitDStreamFF offbits = ZS_BitDStreamFF_init(offbitsBuffer, offbitsSize);
    ZL_DLOG(V9, "OBITS stream size = %u", (unsigned)offbitsSize);
    *bitstream = offbits;
    return ZL_returnSuccess();
}

// decodeOffbits appears to be unused legacy code, kept for potential future use
static ZL_UNUSED_ATTR ZL_Report decodeOffbits(
        uint32_t* offsets,
        uint8_t const* offsetsBuffer,
        size_t numOffsets,
        ZL_RC* in)
{
    if (0) {
        ZS_BitDStreamFF offbits;
        ZL_RET_R_IF_ERR(getBitstream(&offbits, in));
        uint32_t* const offsets1 = offsets;
        for (size_t o = 0; o < numOffsets; o += 3) {
            for (size_t u = 0; u < 3; ++u) {
                uint32_t offset;
                uint8_t const nbBits = offsetsBuffer[o + u];
                offset               = (1u << nbBits);
                offset |= (uint32_t)ZS_BitDStreamFF_read(&offbits, nbBits);
                offsets1[o + u] = offset;
            }
            ZS_BitDStreamFF_reload(&offbits);
        }
        ZL_Report const _report = ZS_BitDStreamFF_finish(&offbits);
        (void)_report;
    } else {
        ZS_BitDStreamFF offbits0;
        ZS_BitDStreamFF offbits1;
        ZL_RET_R_IF_ERR(getBitstream(&offbits0, in));
        ZL_RET_R_IF_ERR(getBitstream(&offbits1, in));
        // ZS_BitDStreamFF offbits2 = getBitstream(in);
        // ZS_BitDStreamFF offbits3 = getBitstream(in);
        for (size_t o = 0; o < numOffsets; o += 4) {
#define X(u, n)                                                                \
    do {                                                                       \
        int const nbBits = offsetsBuffer[o + u * 2 + n];                       \
        uint32_t offset  = 1u << nbBits;                                       \
        offset |= (uint32_t)ZS_BitDStreamFF_read(&offbits##n, (size_t)nbBits); \
        offsets[o + u * 2 + n] = offset;                                       \
    } while (0)
            X(0, 0);
            X(1, 0);
            // X(2, 0);
            ZS_BitDStreamFF_reload(&offbits0);
            X(0, 1);
            X(1, 1);
            // X(2, 1);
            ZS_BitDStreamFF_reload(&offbits1);
#undef X
        }
    }
    return ZL_returnSuccess();
}

#if 0

#    define ZS_DECODE_EXTRA(length)                               \
        do {                                                      \
            length = 0;                                           \
            for (;;) {                                            \
                size_t const _kMaxLength  = (1u << 16) - 1;       \
                size_t const _extraLength = ZL_readLE16(*extras); \
                *extras += 2;                                     \
                length += _extraLength;                           \
                if (_extraLength < _kMaxLength)                   \
                    break;                                        \
            }                                                     \
        } while (0)
ZL_FORCE_INLINE ZL_TARGET_ATTRIBUTE("avx2,bmi2") size_t loopBody(
    uint8_t const* const tokensBuffer,
    uint8_t const** lits,
    uint8_t const** extras,
    uint8_t** op,
    uint32_t const** offs,
    size_t limit,
    bool kDecodeOffbits) {
  (void)kDecodeOffbits;
  size_t i;
  for (i = 0; i < limit; i += 14) {
    bool const fast = isFast(tokensBuffer + i);
    if (ZL_LIKELY(fast)) {
      uint8_t const* tks = tokensBuffer + i;
      uint32_t ofs[8];

      typedef size_t len_t;
#    define X(k)                                          \
        do {                                              \
            uint32_t off                = ofs[k % 7];     \
            len_t const tok             = tks[k];         \
            len_t const llen            = (tok >> 1) & 7; \
            len_t const mlen            = (tok >> 4);     \
            uint8_t* const omatch       = *op + llen;     \
            uint8_t const* const imatch = omatch - off;   \
            uint8_t* const onext        = omatch + mlen;  \
            memcpy(*op, *lits, 8);                        \
            *lits += llen;                                \
            memcpy(omatch, imatch, 16);                   \
            *op = onext;                                  \
        } while (0)

      *offs = getOffsets(ofs, *offs, tokensBuffer + i + 0);
      X(0);
      X(1);
      X(2);
      X(3);
      X(4);
      X(5);
      X(6);
      *offs = getOffsets(ofs, *offs, tokensBuffer + i + 7);
      X(7);
      X(8);
      X(9);
      X(10);
      X(11);
      X(12);
      X(13);
#    undef X
    } else {
      for (size_t t = 0; t < 14; t += 7) {
        uint32_t os[8];
        *offs = getOffsets(os, *offs, tokensBuffer + i + t);
        for (size_t u = 0; u < 7; ++u) {
          uint8_t const token   = tokensBuffer[i + t + u];
          uint32_t const offset = os[u];
          if (ZL_UNLIKELY((token & ~1) == 0)) {
            size_t litlen, matchlen;
            ZS_DECODE_EXTRA(litlen);
            ZS_DECODE_EXTRA(matchlen);
            ZS_wildcopy(*op, *lits, (ptrdiff_t)litlen, ZS_wo_no_overlap);
            *op += litlen;
            *lits += litlen;
            // memcpy(op, op - offset, matchlen);
            uint8_t const* const match = *op - offset;
            // for (size_t x = 0; x < matchlen; ++x) {
            //   op[x] = match[x];
            // }
            ZS_wildcopy(*op, match, (ptrdiff_t)matchlen,
            ZS_wo_src_before_dst); *op += matchlen;
            // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
          } else {
            size_t const litlen   = (token >> 1) & 7;
            size_t const matchlen = (token >> 4);
            // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
            memcpy(*op, *lits, 8);
            *op += litlen;
            *lits += litlen;
            memcpy(*op, *op - offset, 16);
            // uint8_t const* const match = op - offset;
            // for (size_t x = 0; x < matchlen; ++x) {
            //   op[x] = match[x];
            // }
            *op += matchlen;
          }
        }
      }
    }
  }
  return i;
}
#    undef ZS_DECODE_EXTRA

#endif /* #if 0 */

ZL_FORCE_INLINE void
copy(void* dst, void const* src, size_t kStride, size_t kLength)
{
    // memcpy(dst, src, kLength); (void)kStride;
    ZL_ASSERT_EQ(kLength % kStride, 0);
    for (size_t i = 0; i < kLength; i += kStride) {
        memcpy((uint8_t*)dst + i, (uint8_t const*)src + i, kStride);
    }
}

#define ZL_TRY_LET_OR_GOTO(label, var, fn)       \
    ZL_Report var##_res = (fn);                  \
    if (ZL_isError(var##_res)) {                 \
        ZL_DLOG(ERROR, "Error: %s failed", #fn); \
        goto label;                              \
    }                                            \
    size_t const var = ZL_validResult(var##_res);

#define ZL_GOTO_IF(label, cond)                          \
    do {                                                 \
        if (cond) {                                      \
            ZL_DLOG(ERROR, "Error: %s is false", #cond); \
            goto label;                                  \
        }                                                \
    } while (0)

#define ZL_GOTO_IF_ERR(label, res) ZL_GOTO_IF(label, ZL_RES_isError(res))

static ZL_Report ZS_fastDecoder_decompress(
        ZS_decoderCtx* ctx,
        uint8_t* const dst,
        size_t const capacity,
        uint8_t const* const src,
        size_t size)
{
    (void)ctx;
    uint8_t* const ostart = dst;
    uint8_t* op           = ostart;
    uint8_t* const oend   = ostart + capacity;
    uint8_t* const olimit = oend - ZS_WILDCOPY_OVERLENGTH;

    ZL_RC in = ZL_RC_wrap(src, size);

    // Allocations which need to be freed on error
    uint8_t* litsBufferMalloc    = NULL;
    uint16_t* tokensBufferMalloc = NULL;

    // Bounds checks
    ZS_window window;
    uint32_t rep = kMinOffset;
    ZL_RET_R_IF(GENERIC, ZS_window_init(&window, (uint32_t)(oend - ostart), 8));

    ZL_TRY_LET_R(numLiterals, getDecodedSize(in, 1));
    ZL_RET_R_IF_GE(GENERIC, numLiterals, (1 << 30), "too many literals");
    uint8_t const* litsBuffer;
    ZL_TRY_LET_R(rawLits, isRaw(in));
    if (rawLits) {
        litsBuffer = getRawBuffer(&in, 1);
        ZL_RET_R_IF_NULL(GENERIC, litsBuffer);
    } else {
        litsBufferMalloc =
                (uint8_t*)ZL_malloc(numLiterals + ZS_WILDCOPY_OVERLENGTH);
        ZL_RET_R_IF_NULL(allocation, litsBufferMalloc);
        ZL_GOTO_IF_ERR(
                _error,
                decodeLiterals(
                        litsBufferMalloc,
                        numLiterals + ZS_WILDCOPY_OVERLENGTH,
                        &in));
        litsBuffer = litsBufferMalloc;
    }

    uint16_t const* tokensBuffer;
    ZL_TRY_LET_OR_GOTO(_error, numSequences, getDecodedSize(in, 1));
    ZL_TRY_LET_OR_GOTO(_error, rawTokens, isRaw(in));
    ZL_GOTO_IF(_error, numSequences >= (1 << 30));
    if (rawTokens) {
        tokensBuffer = (uint16_t const*)(void const*)getRawBuffer(&in, 2);
        ZL_GOTO_IF(_error, litsBuffer == NULL);
    } else {
        tokensBufferMalloc =
                (uint16_t*)ZL_malloc(numSequences * sizeof(uint16_t));
        ZL_GOTO_IF(_error, tokensBufferMalloc == NULL);
        ZL_GOTO_IF_ERR(
                _error, decodeCodes16(tokensBufferMalloc, numSequences, &in));
        tokensBuffer = tokensBufferMalloc;
    }

    uint64_t offsetsSize;
    {
        ZL_RESULT_OF(uint64_t) res = ZL_RC_popVarint(&in);
        ZL_GOTO_IF_ERR(_error, res);
        offsetsSize = ZL_RES_value(res);
    }
    uint8_t const* offsets = ZL_RC_ptr(&in);
    ZL_GOTO_IF(_error, offsetsSize > ZL_RC_avail(&in));
    ZL_RC_advance(&in, offsetsSize);

    uint64_t extrasSize;
    {
        ZL_RESULT_OF(uint64_t) res = ZL_RC_popVarint(&in);
        ZL_GOTO_IF_ERR(_error, res);
        extrasSize = ZL_RES_value(res);
    }
    uint8_t const* const extrasBuffer = ZL_RC_ptr(&in);
    ZL_GOTO_IF(_error, extrasSize > ZL_RC_avail(&in));
    ZL_RC_advance(&in, extrasSize);
    ZL_DLOG(TRANSFORM, "Found %u extra bytes", (unsigned)extrasSize);

    ZL_GOTO_IF(_error, ZL_RC_avail(&in) != 0);

    ZS_window_update(&window, ostart, (size_t)(oend - ostart));
    uint8_t const* lits   = litsBuffer;
    uint8_t const* extras = extrasBuffer;

    uint8_t const* const litsEnd    = lits + numLiterals;
    uint8_t const* const extrasEnd  = extras + extrasSize;
    uint8_t const* const offsetsEnd = offsets + offsetsSize;

    uint8_t const* const litsLimit = litsEnd - ZS_WILDCOPY_OVERLENGTH;

#define ZS_DECODE_EXTRA(length)                                  \
    do {                                                         \
        ZL_GOTO_IF(_error, (extrasEnd - extras) < 1);            \
        length = *extras++;                                      \
        if (ZL_UNLIKELY(length == 255))                          \
            for (;;) {                                           \
                ZL_GOTO_IF(_error, (extrasEnd - extras) < 2);    \
                size_t const _kMaxLength  = (1u << 16) - 1;      \
                size_t const _extraLength = ZL_readLE16(extras); \
                extras += 2;                                     \
                length += _extraLength;                          \
                if (_extraLength < _kMaxLength)                  \
                    break;                                       \
            }                                                    \
    } while (0)

    size_t const limit = numSequences - 32;

    size_t i = 0;
    // TODO: We have a bad case where the first sequence is 15 literals, since
    // then we will copy a match from offset 16. This is clearly bad, but this
    // code is not currently in use, so just special case it.
    if (numSequences > 0 && ZL_read16(&tokensBuffer[0]) == 60) {
        size_t const llen =
                (ZL_read16(&tokensBuffer[0]) >> kTokenOFBits) & kTokenLLMask;
        ZL_GOTO_IF(_error, llen > numLiterals);
        ZL_GOTO_IF(_error, llen > (size_t)(oend - op));
        memcpy(op, lits, llen);
        op += llen;
        lits += llen;
        ++i;
    }

    if (1) {
        size_t const fastTotalML     = 16 * (1u << kTokenMLBits);
        size_t const fastTotalLL     = 16 * (1u << kTokenLLBits);
        size_t const fastTotalOut    = fastTotalLL + fastTotalML;
        size_t const fastOffsetBytes = 16 * 3;
        if (numSequences >= 32 && op < olimit && lits < litsLimit) {
            for (; i < limit && op + fastTotalOut <= olimit
                 && lits + fastTotalLL <= litsLimit
                 && (size_t)(offsetsEnd - offsets) >= fastOffsetBytes;
                 i += 16) {
                bool const fast = isFast(tokensBuffer + i);
                if (ZL_LIKELY(fast)) {
                    uint16_t const* tks = tokensBuffer + i;
                    uint32_t ofs[4];

                    typedef size_t len_t;
#define X(k)                                                           \
    do {                                                               \
        uint32_t off    = ofs[k & 3];                                  \
        len_t const tok = ZL_read16(&tks[k]);                          \
        ZL_ASSERT_NE(tok & (size_t)~3, 0);                             \
        len_t const llen = (tok >> kTokenOFBits) & kTokenLLMask;       \
        len_t const mlen =                                             \
                (tok >> (kTokenOFBits + kTokenLLBits)) & kTokenMLMask; \
        uint8_t* const omatch       = op + llen;                       \
        uint8_t const* const imatch = omatch - off;                    \
        uint8_t* const onext        = omatch + mlen;                   \
        memcpy(op, lits, 1u << kTokenLLBits);                          \
        lits += llen;                                                  \
        ZL_GOTO_IF(_error, (off) > (size_t)(omatch - ostart));         \
        ZL_GOTO_IF(_error, (off) < kMinOffset);                        \
        copy(omatch, imatch, kMinOffset, 1u << kTokenMLBits);          \
        op = onext;                                                    \
    } while (0)

                    offsets = getOffsets4(ofs, offsets, tks + 0, rep);
                    X(0);
                    X(1);
                    X(2);
                    X(3);
                    offsets = getOffsets4(ofs, offsets, tks + 4, ofs[3]);
                    X(4);
                    X(5);
                    X(6);
                    X(7);
                    offsets = getOffsets4(ofs, offsets, tks + 8, ofs[3]);
                    X(8);
                    X(9);
                    X(10);
                    X(11);
                    offsets = getOffsets4(ofs, offsets, tks + 12, ofs[3]);
                    X(12);
                    X(13);
                    X(14);
                    X(15);
                    rep = ofs[3];
                } else {
                    for (size_t t = 0; t < 16; t += 4) {
                        uint32_t os[4];
                        offsets = getOffsets4(
                                os, offsets, tokensBuffer + i + t, rep);
                        for (size_t u = 0; u < 4; ++u) {
                            uint16_t const token =
                                    ZL_read16(&tokensBuffer[i + t + u]);
                            uint32_t const offset = os[u];
                            size_t litlen, matchlen;
                            if (ZL_UNLIKELY((token & ~3) == 0)) {
                                ZS_DECODE_EXTRA(litlen);
                                ZS_DECODE_EXTRA(matchlen);
                                if (ZL_LIKELY(
                                            lits + litlen <= litsLimit
                                            && op + litlen + matchlen
                                                    <= olimit)) {
                                    ZS_wildcopy(
                                            op,
                                            lits,
                                            (ptrdiff_t)litlen,
                                            ZS_wo_no_overlap);
                                    op += litlen;
                                    lits += litlen;
                                    ZL_GOTO_IF(_error, offset < kMinOffset);
                                    ZL_GOTO_IF(
                                            _error,
                                            offset > (size_t)(op - ostart));
                                    // memcpy(op, op - offset, matchlen);
                                    uint8_t const* const match = op - offset;
                                    // for (size_t x = 0; x < matchlen; ++x) {
                                    //   op[x] = match[x];
                                    // }
                                    ZS_wildcopy(
                                            op,
                                            match,
                                            (ptrdiff_t)matchlen,
                                            ZS_wo_src_before_dst);
                                    op += matchlen;
                                    // ZL_DLOG(V9, "llen = %zu | mlen = %zu",
                                    // litlen, matchlen);
                                    continue;
                                }
                            } else {
                                litlen = (token >> kTokenOFBits) & kTokenLLMask;
                                matchlen =
                                        (token >> (kTokenOFBits + kTokenLLBits))
                                        & kTokenMLMask;
                                if (ZL_LIKELY(
                                            lits + kTokenLLCopy <= litsLimit
                                            && op + kTokenLLCopy + kTokenMLCopy
                                                    <= olimit)) {
                                    // ZL_DLOG(V9, "llen = %zu | mlen = %zu",
                                    // litlen, matchlen);
                                    memcpy(op, lits, kTokenLLCopy);
                                    op += litlen;
                                    lits += litlen;
                                    ZL_GOTO_IF(_error, offset < kMinOffset);
                                    ZL_GOTO_IF(
                                            _error,
                                            offset > (size_t)(op - ostart));
                                    copy(op,
                                         op - offset,
                                         kMinOffset,
                                         kTokenMLCopy);
                                    // uint8_t const* const match = op - offset;
                                    // for (size_t x = 0; x < matchlen; ++x) {
                                    //   op[x] = match[x];
                                    // }
                                    op += matchlen;
                                    continue;
                                }
                            }
                            ZL_GOTO_IF(
                                    _error, litlen > (size_t)(litsEnd - lits));
                            ZL_GOTO_IF(
                                    _error,
                                    litlen + matchlen > (size_t)(oend - op));
                            ZS_safecopy(
                                    op,
                                    lits,
                                    (ptrdiff_t)litlen,
                                    ZS_wo_no_overlap);
                            op += litlen;
                            lits += litlen;
                            ZL_GOTO_IF(_error, offset < kMinOffset);
                            ZL_GOTO_IF(_error, offset > (size_t)(op - ostart));
                            // memcpy(op, op - offset, matchlen);
                            uint8_t const* const match = op - offset;
                            // for (size_t x = 0; x < matchlen; ++x) {
                            //   op[x] = match[x];
                            // }
                            ZS_safecopy(
                                    op,
                                    match,
                                    (ptrdiff_t)matchlen,
                                    ZS_wo_src_before_dst);
                            op += matchlen;
                            // ZL_DLOG(V9, "llen = %zu | mlen = %zu",
                            // litlen, matchlen);
                        }
                        rep = os[3];
                    }
                }
            }
        }
        for (; i < numSequences; ++i) {
            uint16_t const token = ZL_read16(&tokensBuffer[i]);
            uint32_t offset;
            int const ocode = token & 3;
            if (ocode == 0) {
                offset = rep;
            } else if (ocode == 1) {
                ZL_GOTO_IF(_error, (size_t)(offsetsEnd - offsets) < 1);
                offset = *offsets++;
                rep    = offset;
            } else if (ocode == 2) {
                ZL_GOTO_IF(_error, (size_t)(offsetsEnd - offsets) < 2);
                offset = ZL_readLE16(offsets);
                offsets += 2;
                rep = offset;
            } else {
                ZL_GOTO_IF(_error, (size_t)(offsetsEnd - offsets) < 3);
                offset = ZL_readLE24(offsets);
                offsets += 3;
                rep = offset;
            }
            if (ZL_UNLIKELY((token & ~3) == 0)) {
                size_t litlen, matchlen;
                ZS_DECODE_EXTRA(litlen);
                ZS_DECODE_EXTRA(matchlen);
                ZL_GOTO_IF(_error, litlen > (size_t)(litsEnd - lits));
                ZL_GOTO_IF(_error, litlen + matchlen > (size_t)(oend - op));
                ZS_safecopy(op, lits, (ptrdiff_t)litlen, ZS_wo_no_overlap);
                op += litlen;
                ZL_GOTO_IF(_error, offset < kMinOffset);
                ZL_GOTO_IF(_error, offset > (size_t)(op - ostart));
                lits += litlen;
                uint8_t const* const match = op - offset;
                ZS_safecopy(
                        op, match, (ptrdiff_t)matchlen, ZS_wo_src_before_dst);
                op += matchlen;
                // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
            } else {
                size_t const litlen = (token >> kTokenOFBits) & kTokenLLMask;
                size_t const matchlen =
                        (token >> (kTokenOFBits + kTokenLLBits)) & kTokenMLMask;
                ZL_GOTO_IF(_error, litlen > (size_t)(litsEnd - lits));
                ZL_GOTO_IF(_error, litlen + matchlen > (size_t)(oend - op));
                // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
                ZS_safecopy(op, lits, (ptrdiff_t)litlen, ZS_wo_no_overlap);
                op += litlen;
                ZL_GOTO_IF(_error, offset < kMinOffset);
                ZL_GOTO_IF(_error, offset > (size_t)(op - ostart));
                lits += litlen;
                ZS_safecopy(
                        op,
                        op - offset,
                        (ptrdiff_t)matchlen,
                        ZS_wo_src_before_dst);
                // uint8_t const* const match = op - offset;
                // for (size_t x = 0; x < matchlen; ++x) {
                //   op[x] = match[x];
                // }
                op += matchlen;
            }
        }
    } else {
#if 0
        uint8_t const* offs = offset16 - 2;
        if (numSequences > 0)
          ZL_REQUIRE_NE(tokensBuffer[0] & 1, 0);
        if (0)
          i = 0;
        else if (numSequences >= 16)
          for (i = 0; i < limit; i += 14) {
            bool const fast = isFast(tokensBuffer + i);
            if (ZL_LIKELY(fast)) {
              uint8_t const* tks = tokensBuffer + i;
              uint16_t ofs[16];

              typedef size_t len_t;
#    define X(k)                                          \
        do {                                              \
            uint32_t off                = ofs[k % 7];     \
            len_t const tok             = tks[k];         \
            len_t const llen            = (tok >> 1) & 7; \
            len_t const mlen            = (tok >> 4);     \
            uint8_t* const omatch       = op + llen;      \
            uint8_t const* const imatch = omatch - off;   \
            uint8_t* const onext        = omatch + mlen;  \
            memcpy(op, lits, 8);                          \
            lits += llen;                                 \
            memcpy(omatch, imatch, 16);                   \
            op = onext;                                   \
        } while (0)

              offs = getOffsets16(ofs, offs, tokensBuffer + i + 0);
              X(0);
              X(1);
              X(2);
              X(3);
              X(4);
              X(5);
              X(6);
              offs = getOffsets16(ofs, offs, tokensBuffer + i + 7);
              X(7);
              X(8);
              X(9);
              X(10);
              X(11);
              X(12);
              X(13);
            } else {
              for (size_t t = 0; t < 14; t += 7) {
                uint16_t os[8];
                offs = getOffsets16(os, offs, tokensBuffer + i + t);
                for (size_t u = 0; u < 7; ++u) {
                  uint8_t const token   = tokensBuffer[i + t + u];
                  uint32_t const offset = os[u];
                  if (ZL_UNLIKELY((token & ~1) == 0)) {
                    size_t litlen, matchlen;
                    ZS_DECODE_EXTRA(litlen);
                    ZS_DECODE_EXTRA(matchlen);
                    ZS_wildcopy(op, lits, (ptrdiff_t)litlen,
                    ZS_wo_no_overlap); op += litlen; lits += litlen;
                    // memcpy(op, op - offset, matchlen);
                    uint8_t const* const match = op - offset;
                    // for (size_t x = 0; x < matchlen; ++x) {
                    //   op[x] = match[x];
                    // }
                    ZS_wildcopy(
                        op, match, (ptrdiff_t)matchlen,
                        ZS_wo_src_before_dst);
                    op += matchlen;
                    // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen,
                    matchlen);
                  } else {
                    size_t const litlen   = (token >> 1) & 7;
                    size_t const matchlen = (token >> 4);
                    // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen,
                    matchlen); memcpy(op, lits, 8); op += litlen; lits +=
                    litlen; memcpy(op, op - offset, 16);
                    // uint8_t const* const match = op - offset;
                    // for (size_t x = 0; x < matchlen; ++x) {
                    //   op[x] = match[x];
                    // }
                    op += matchlen;
                  }
                }
              }
            }
          }
        else
          i = 0;
        rep = ZL_readLE16(offs);
        offs += 2;
        for (; i < numSequences; ++i) {
          uint8_t const token = tokensBuffer[i];
          uint32_t offset;
          if ((token & 1) == 0) {
            offset = rep;
          } else {
            offset = ZL_readLE16(offs);
            offs += 2;
            rep = offset;
          }
          if (ZL_UNLIKELY((token & ~1) == 0)) {
            size_t litlen, matchlen;
            ZS_DECODE_EXTRA(litlen);
            ZS_DECODE_EXTRA(matchlen);
            ZS_wildcopy(op, lits, (ptrdiff_t)litlen, ZS_wo_no_overlap);
            op += litlen;
            lits += litlen;
            // memcpy(op, op - offset, matchlen);
            uint8_t const* const match = op - offset;
            // for (size_t x = 0; x < matchlen; ++x) {
            //   op[x] = match[x];
            // }
            ZS_wildcopy(op, match, (ptrdiff_t)matchlen,
            ZS_wo_src_before_dst); op += matchlen;
            // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
          } else {
            size_t const litlen   = (token >> 1) & 7;
            size_t const matchlen = (token >> 4);
            // ZL_DLOG(V9, "llen = %zu | mlen = %zu", litlen, matchlen);
            memcpy(op, lits, 8);
            op += litlen;
            lits += litlen;
            memcpy(op, op - offset, 16);
            // uint8_t const* const match = op - offset;
            // for (size_t x = 0; x < matchlen; ++x) {
            //   op[x] = match[x];
            // }
            op += matchlen;
          }
        }
#endif /* #if 0 */
    }

    size_t const lastlits = (size_t)((litsBuffer + numLiterals) - lits);
    if (lastlits > (size_t)(oend - op))
        goto _error;
    memcpy(op, lits, lastlits);
    op += lastlits;

#undef ZS_DECODE_EXTRA

    ZL_ASSERT(op <= oend);

    ZL_free(tokensBufferMalloc);
    ZL_free(litsBufferMalloc);

    return ZL_returnValue((size_t)(op - ostart));

_error:
    ZL_free(tokensBufferMalloc);
    ZL_free(litsBufferMalloc);
    ZL_RET_R_ERR(GENERIC);
}

const ZS_decoder ZS_fastDecoder = {
    .name        = "fast",
    .ctx_create  = ZS_decoderCtx_create,
    .ctx_release = ZS_decoderCtx_release,
    .ctx_reset   = ZS_decoderCtx_reset,
    .decompress  = ZS_fastDecoder_decompress,
};
